typedef signed   char  int8_t;
typedef unsigned char  uint8_t;
typedef signed   short int16_t;
typedef unsigned short uint16_t;
typedef signed   int   int32_t;
typedef unsigned int   uint32_t;

typedef uint8_t        pel_t;
typedef uint16_t       lpel_t;

/* ---------------------------------------------------------------------------
 * interpolate filter (luma) */
__constant const int8_t INTPL_FILTERS[4][8] = {
    {  0, 0,   0, 64,  0,  0,  0,  0 }, /* for full-pixel, no use */
    { -1, 4, -10, 57, 19, -7,  3, -1 },
    { -1, 4, -11, 40, 40, -11, 4, -1 },
    { -1, 3,  -7, 19, 57, -10, 4, -1 }
};
__constant const int INTPL_SHIFT = 6;
__constant const int INTPL_OFFSET = (1 << 5);

#define MIN(a, b)                      ((a) < (b) ? (a) : (b))
#define MAX(a, b)                      ((a) > (b) ? (a) : (b))
#define CLIP3(min_val, max_val, val)   MAX(min_val, MIN(max_val, val))

__kernel void cl_kernel_mul_add(__global const pel_t *src_a,
                                __global uint8_t     *result_p,
                                const int offset_p2,
                                const int offset_p3,
                                const int offset_t1,
                                const int offset_t2,
                                const int offset_t3,
                                const int i_stride)
{
    __global pel_t  *result_p1 = (__global pel_t  *)(result_p + 0);
    __global pel_t  *result_p2 = (__global pel_t  *)(result_p + offset_p2);
    __global pel_t  *result_p3 = (__global pel_t  *)(result_p + offset_p3);
    __global lpel_t *result_t1 = (__global lpel_t *)(result_p + offset_t1);
    __global lpel_t *result_t2 = (__global lpel_t *)(result_p + offset_t2);
    __global lpel_t *result_t3 = (__global lpel_t *)(result_p + offset_t3);
    const int min_x = 0;
    const int max_x = i_stride - 1;
    int gid_x = get_global_id(0);
    int gid_y = get_global_id(1);
    int pos = gid_y * i_stride;

    int j0 = CLIP3(min_x, max_x, gid_x - 3);
    int j1 = CLIP3(min_x, max_x, gid_x - 2);
    int j2 = CLIP3(min_x, max_x, gid_x - 1);
    int j3 = CLIP3(min_x, max_x, gid_x - 0);
    int j4 = CLIP3(min_x, max_x, gid_x + 1);
    int j5 = CLIP3(min_x, max_x, gid_x + 2);
    int j6 = CLIP3(min_x, max_x, gid_x + 3);
    int j7 = CLIP3(min_x, max_x, gid_x + 4);
    int sum;
    /* position 1 */
    sum = INTPL_FILTERS[1][0] * src_a[pos + j0]
        + INTPL_FILTERS[1][1] * src_a[pos + j1]
        + INTPL_FILTERS[1][2] * src_a[pos + j2]
        + INTPL_FILTERS[1][3] * src_a[pos + j3]
        + INTPL_FILTERS[1][4] * src_a[pos + j4]
        + INTPL_FILTERS[1][5] * src_a[pos + j5]
        + INTPL_FILTERS[1][6] * src_a[pos + j6]
        + INTPL_FILTERS[1][7] * src_a[pos + j7];
    result_t1[pos + gid_x] = sum;
    sum = (sum + INTPL_OFFSET) >> INTPL_SHIFT;
    result_p1[pos + gid_x] = CLIP3(0, 255, sum);

    /* position 2 */
    sum = INTPL_FILTERS[2][0] * src_a[pos + j0]
        + INTPL_FILTERS[2][1] * src_a[pos + j1]
        + INTPL_FILTERS[2][2] * src_a[pos + j2]
        + INTPL_FILTERS[2][3] * src_a[pos + j3]
        + INTPL_FILTERS[2][4] * src_a[pos + j4]
        + INTPL_FILTERS[2][5] * src_a[pos + j5]
        + INTPL_FILTERS[2][6] * src_a[pos + j6]
        + INTPL_FILTERS[2][7] * src_a[pos + j7];
    result_t2[pos + gid_x] = sum;
    sum = (sum + INTPL_OFFSET) >> INTPL_SHIFT;
    result_p2[pos + gid_x] = CLIP3(0, 255, sum);

    /* position 3 */
    sum = INTPL_FILTERS[3][0] * src_a[pos + j0]
        + INTPL_FILTERS[3][1] * src_a[pos + j1]
        + INTPL_FILTERS[3][2] * src_a[pos + j2]
        + INTPL_FILTERS[3][3] * src_a[pos + j3]
        + INTPL_FILTERS[3][4] * src_a[pos + j4]
        + INTPL_FILTERS[3][5] * src_a[pos + j5]
        + INTPL_FILTERS[3][6] * src_a[pos + j6]
        + INTPL_FILTERS[3][7] * src_a[pos + j7];
    result_t3[pos + gid_x] = sum;
    sum = (sum + INTPL_OFFSET) >> INTPL_SHIFT;
    result_p3[pos + gid_x] = CLIP3(0, 255, sum);
}
